#include "CudaRGB32ToYUV444Service.h"

__global__ void RGB32ToYUV444(uchar *RGB, uchar *y, uchar *u, uchar *v, int w, int h)
{
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if(tid<w*h){
        y[tid] = 0.299*(int)RGB[3*tid] + 0.587*(int)RGB[3*tid+1] + 0.114*(int)RGB[3*tid+2];
        u[tid] = -0.1684*(int)RGB[3*tid] - 0.3316*(int)RGB[3*tid+1] +0.5*(int)RGB[3*tid+2] + 128;
        v[tid] = 0.5*(int)RGB[3*tid] - 0.4187*(int)RGB[3*tid+1] -0.0813*(int)RGB[3*tid+2] + 128;
        u[tid] = (u[tid] < 0) ? 0 : ((u[tid] > 255) ? 255 : u[tid]);
        v[tid] = (v[tid] < 0) ? 0 : ((v[tid] > 255) ? 255 : v[tid]);
    }
}
__global__ void RGB32ToYUV420(uchar *RGB, uchar *y, uchar *u, uchar *v, int w, int h){
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if(tid<w * h){
        y[tid] = 0.299*(int)RGB[3*tid] + 0.587*(int)RGB[3*tid+1] + 0.114*(int)RGB[3*tid+2];
        y[tid] = (y[tid] < 0) ? 0 : ((y[tid] > 255) ? 255 : y[tid]);
        if((tid/w)%2 == 0 && (tid%w)%2 == 0){
            int index = ((tid%w)/2) + ((tid/w)/2)*(w/2);
            u[index] = -0.1684*(int)RGB[3*tid] - 0.3316*(int)RGB[3*tid+1] +0.5*(int)RGB[3*tid+2] + 128;
            v[index] = 0.5*(int)RGB[3*tid] - 0.4187*(int)RGB[3*tid+1] -0.0813*(int)RGB[3*tid+2] + 128;
            u[index] = (u[index] < 0) ? 0 : ((u[index] > 255) ? 255 : u[index]);
            v[index] = (v[index] < 0) ? 0 : ((v[index] > 255) ? 255 : v[index]);
        }
    }
}

CudaRGB32ToYUV444Service::CudaRGB32ToYUV444Service()
{
}

CudaRGB32ToYUV444Service::CudaRGB32ToYUV444Service(int w, int h)
{
    //分配空间
    this->w = w;
    this->h = h;
    cudaMalloc((void**)&d_RGB,3*w*h*sizeof(uchar));
    cudaMalloc((void**)&d_y,w*h*sizeof(uchar));
    cudaMalloc((void**)&d_u,w*h*sizeof(uchar));
    cudaMalloc((void**)&d_v,w*h*sizeof(uchar));
}

CudaRGB32ToYUV444Service::~CudaRGB32ToYUV444Service()
{
    //释放
    cudaFree(d_RGB);
    cudaFree(d_y);
    cudaFree(d_u);
    cudaFree(d_v);
}

void CudaRGB32ToYUV444Service::RGB32ToYUV444Interface(uchar *RGB, uchar *y, uchar *u, uchar *v, int w, int h)
{
    //拷贝到gpu
    cudaMemcpy( d_RGB, RGB, 3*w*h*sizeof(uchar), cudaMemcpyHostToDevice);
    //配置gpu线程模型
    int blockSize = 256;
    int gridSize = (w*h + blockSize - 1) / blockSize;
    RGB32ToYUV444<<<gridSize, blockSize>>>(d_RGB,d_y,d_u,d_v,w,h);
    cudaMemcpy( y, d_y, w*h*sizeof(uchar), cudaMemcpyDeviceToHost);
    cudaMemcpy( u, d_u, w*h*sizeof(uchar), cudaMemcpyDeviceToHost);
    cudaMemcpy( v, d_v, w*h*sizeof(uchar), cudaMemcpyDeviceToHost);
}
void CudaRGB32ToYUV444Service::RGB32ToYUV420Interface(uchar *RGB, uchar *y, uchar *u, uchar *v, int w, int h)
{
    //拷贝到gpu
    cudaMemcpy( d_RGB, RGB, 3*w*h*sizeof(uchar), cudaMemcpyHostToDevice);
    //配置gpu线程模型
    int blockSize = 256;
    int gridSize = (w*h + blockSize - 1) / blockSize;
    RGB32ToYUV420<<<gridSize, blockSize>>>(d_RGB,d_y,d_u,d_v,w,h);
    cudaMemcpy( y, d_y, w*h*sizeof(uchar), cudaMemcpyDeviceToHost);
    cudaMemcpy( u, d_u,0.25* w*h*sizeof(uchar), cudaMemcpyDeviceToHost);
    cudaMemcpy( v, d_v, 0.25*w*h*sizeof(uchar), cudaMemcpyDeviceToHost);
}

